/*
 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
 *
 * NVIDIA Corporation and its licensors retain all intellectual property and 
 * proprietary rights in and to this software and related documentation. 
 * Any use, reproduction, disclosure, or distribution of this software 
 * and related documentation without an express license agreement from
 * NVIDIA Corporation is strictly prohibited.
 *
 * Please refer to the applicable NVIDIA end user license agreement (EULA) 
 * associated with this source code for terms and conditions that govern 
 * your use of this NVIDIA software.
 * 
 */



#ifndef QUASIRANDOMGENERATOR_KERNEL_CUH
#define QUASIRANDOMGENERATOR_KERNEL_CUH


#include <stdio.h>
#include <stdlib.h>
#include <cutil_inline.h>
#include "realtype.h"
#include "quasirandomGenerator_common.h"



//Fast integer multiplication
#define MUL(a, b) __umul24(a, b)



////////////////////////////////////////////////////////////////////////////////
// Niederreiter quasirandom number generation kernel
////////////////////////////////////////////////////////////////////////////////
static __constant__ unsigned int c_Table[QRNG_DIMENSIONS][QRNG_RESOLUTION];


////////////////////////////////////////////////////////////////////////////////
// Aqui comienza el TSP
////////////////////////////////////////////////////////////////////////////////
__device__ int calcularCosto(int N, int *w, int *distancias, int thread){
  int costo = 0;

  for (int i=0; i<N-1; i++){
		int n1 = w[thread * N + i];
    int n2 = w[thread * N + i + 1];
    costo += distancias[n1*N+n2];
  }
  int n1 = w[thread * N];
  int n2 = w[thread * N + N - 1];
  costo += distancias[n1*N+n2];

  return costo;
}

static __global__ void tspKernel(int NMax, int* distancias, int optimo, int* w, int* wPrima){
		unsigned int tid = MUL(blockDim.x, blockIdx.x) + threadIdx.x;

		if(tid<NMax) w[tid] = 1;
}
static __global__ void quasirandomGeneratorKernel(float *d_Output, unsigned int seed, unsigned int N){
    unsigned int *dimBase = &c_Table[threadIdx.y][0];
    unsigned int      tid = MUL(blockDim.x, blockIdx.x) + threadIdx.x;
    unsigned int  threadN = MUL(blockDim.x, gridDim.x);
}

//Host-side interface
static void tspGPU(int NMax, int* distancias_d, int optimo, int *solucion_h, int* w_d, int* wPrima_d){
  /*dim3 threads(128, QRNG_DIMENSIONS);
  
	cudaMemset(w_d, 0, NMax * sizeof(int));

	int nThreads = NMax;
	printf("lanzando %d threads en 1 grid en kernel ...\n", nThreads);
	tspKernel<<<128, threads>>>(NMax, distancias_d, optimo, w_d, wPrima_d);
  cutilCheckMsg("tspKernel() execution failed.\n");
	printf("kernel finalizado.\n");

	cudaThreadSynchronize();

	//copio el resultado de la memoria del device al host
	cudaMemcpy(solucion_h, w_d, NMax * sizeof(int), cudaMemcpyDeviceToHost);

	printf("i = 0 : %d\n",solucion_h[0]);*/
}

////////////////////////////////////////////////////////////////////////////////
// De aqui en adelante queda todo lo relacionado con el programa <QuasirandomGenerator>
////////////////////////////////////////////////////////////////////////////////
__device__ float cudaRandom(unsigned int seed, unsigned int N){
		float resultado = 0;
    unsigned int *dimBase = &c_Table[threadIdx.y][0];
    unsigned int      tid = MUL(blockDim.x, blockIdx.x) + threadIdx.x;
    unsigned int  threadN = MUL(blockDim.x, gridDim.x);

    for(unsigned int pos = tid; pos < N; pos += threadN){
        unsigned int result = 0;
        unsigned int data = seed + pos;

        for(int bit = 0; bit < QRNG_RESOLUTION; bit++, data >>= 1)
            if(data & 1) result ^= dimBase[bit];

        resultado = (float)(result + 1) * INT_SCALE;
    }
		return resultado;
}

//Table initialization routine
static void initTableGPU(unsigned int tableCPU[QRNG_DIMENSIONS][QRNG_RESOLUTION]){
    cutilSafeCall( cudaMemcpyToSymbol(
        c_Table,
        tableCPU,
        QRNG_DIMENSIONS * QRNG_RESOLUTION * sizeof(unsigned int)
    ) );
}

//Host-side interface
static void quasirandomGeneratorGPU(float *d_Output, unsigned int seed, unsigned int N){
    dim3 threads(128, QRNG_DIMENSIONS);
    quasirandomGeneratorKernel<<<128, threads>>>(d_Output, seed, N);
    cutilCheckMsg("quasirandomGeneratorKernel() execution failed.\n");
}

#endif
